-
Notifications
You must be signed in to change notification settings - Fork 569
[WIP] Remove legacy FIL #6603
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
base: branch-25.06
Are you sure you want to change the base?
[WIP] Remove legacy FIL #6603
Conversation
hcho3
commented
Apr 29, 2025
- Remove legacy FIL from gtest
- Remove legacy FIL from the Python layer of cuML RF
Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually. Contributors can view more details about this message here. |
// Allocate workspace and perform segmented reduce | ||
thrust::device_vector<kv_type> workspace2( | ||
params.n_rows + temp_storage_bytes / sizeof(kv_type) + 1); | ||
cub::DeviceSegmentedReduce::ArgMax( | ||
thrust::raw_pointer_cast(workspace2.data() + params.n_rows), temp_storage_bytes, | ||
workspace->begin(), workspace2.begin(), | ||
params.n_rows, offsets_it, offsets_it + 1); | ||
thrust::transform(workspace2.begin(), workspace2.begin() + params.n_rows, pred->begin(), | ||
[] __device__ (kv_type x) -> int { return x.key; }); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Note. The new FIL always returns probabilities for classifiers. If class outputs are desired, we need to perform argmax.
// Handle a degenerate tree with a childless root node | ||
if (aligned_data.inner_data.distant_offset == 0) { | ||
return offset_type{} + aligned_data.inner_data.distant_offset; | ||
} |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
This provides a fix for a subtle memory bug. Given a degenerate tree with a single childless root node, the node struct is initialized with distant_offset = 0
. When the inference kernel accesses the next node at index condition + (distant_offset - 1)
, it tries to access index -1
, leading to out-of-bounds access.
I was able to discover the bug with the help of property-based gtests.